// #ifndef BOUNDARY_GPU_CUH
// #define BOUNDARY_GPU_CUH

// #include "Global_Variables_extern_gpu.cuh"

// #endif
/* boundary condition */
#define BEFORE_COLLSION (1)
#define AFTER_COLLSION (2)
/* Select Non-equilibrium extrapolation method: BEFORE_COLLSION / AFTER_COLLSION */
#define GUO_EXTRAPOLATION (AFTER_COLLSION)

#define NONE_NRBC (0) 
#define ABC (1)
// Select FLAG_NRBC method: NONE_NRBC / ABC  (which means none NRBC / absorbing boundary conditions)
#define FLAG_NRBC (NONE_NRBC) 

#define BOUNDARY_PERIODIC (0)  //Periodic boundary
#define BOUNDARY_BB_OFFWALL (1)  //off-wall bounce back (the first Internal grid point is 0.5*dx)
#define BOUNDARY_NEEM (2)  //Zhao-Li Guo, et.al., 2002, Non-equilibrium extrapolation method
#define BOUNDARY_ZHOUHE (3)  //Zhou and He (It hasn't been implemented yet)
#define BOUNDARY_ABC (4)  //Guo's Non-equilibrium extrapolation method couple absorption boundary conditions

#define FLAG_BOUNDARY_X0 (BOUNDARY_BB_OFFWALL)  // Specify the first boundary in the x-direction, i.e., the back boundary
#define FLAG_BOUNDARY_X1 (BOUNDARY_BB_OFFWALL)  // Specify the second boundary in the x-direction, i.e., the front boundary
#define FLAG_BOUNDARY_Y0 (BOUNDARY_BB_OFFWALL)  // Specify the first boundary in the y-direction, i.e., the left boundary
#define FLAG_BOUNDARY_Y1 (BOUNDARY_BB_OFFWALL)  // Specify the second boundary in the y-direction, i.e., the right boundary
#define FLAG_BOUNDARY_Z0 (BOUNDARY_BB_OFFWALL)  // Specify the first boundary in the z-direction, i.e., the bottom boundary
#define FLAG_BOUNDARY_Z1 (BOUNDARY_BB_OFFWALL)  // Specify the second boundary in the z-direction, i.e., the top boundary

extern __global__ void wall_BB_BC_GPU_face_xy(I_INT imax, T_P *pdf_old_gpu, I_INT pitch, int ntime);
extern __global__ void wall_BB_BC_GPU_face_yz(I_INT imax, T_P *pdf_old_gpu, I_INT pitch, int ntime);
extern __global__ void wall_BB_BC_GPU_face_xz(I_INT imax, T_P *pdf_old_gpu, I_INT pitch, int ntime);

extern __global__ void periodic_BC_GPU_face_xy(I_INT imax, T_P *pdf_old_gpu, I_INT pitch, int ntime);
extern __global__ void periodic_BC_GPU_face_yz(I_INT imax, T_P *pdf_old_gpu, I_INT pitch, int ntime);
extern __global__ void periodic_BC_GPU_face_xz(I_INT imax, T_P *pdf_old_gpu, I_INT pitch, int ntime);
extern __global__ void periodic_BC_GPU_edge_y(I_INT imax, T_P *pdf_old_gpu, I_INT pitch, int ntime);

extern __global__ void BC_Guo2002_GPU_face_xz_save(I_INT imax, T_P *pdf_old_gpu, I_INT pitch, int ntime, T_P *Boundary_xz0_gpu, T_P *Boundary_xz1_gpu);
extern __global__ void BC_Guo2002_GPU_face_xz_load(I_INT imax, T_P *pdf_old_gpu, I_INT pitch, int ntime, T_P *Boundary_xz0_gpu, T_P *Boundary_xz1_gpu);
extern __global__ void BC_Guo2002_GPU_face_xy_save(I_INT imax, T_P *pdf_old_gpu, I_INT pitch, int ntime, T_P *Boundary_xy0_gpu, T_P *Boundary_xy1_gpu);
extern __global__ void BC_Guo2002_GPU_face_xy_load(I_INT imax, T_P *pdf_old_gpu, I_INT pitch, int ntime, T_P *Boundary_xy0_gpu, T_P *Boundary_xy1_gpu);
extern __global__ void BC_Guo2002_GPU_face_yz_save(I_INT imax, T_P *pdf_old_gpu, I_INT pitch, int ntime, T_P *Boundary_yz0_gpu, T_P *Boundary_yz1_gpu);
extern __global__ void BC_Guo2002_GPU_face_yz_load(I_INT imax, T_P *pdf_old_gpu, I_INT pitch, int ntime, T_P *Boundary_yz0_gpu, T_P *Boundary_yz1_gpu);
extern __global__ void BC_NEEM_AfterColl_GPU_face_xz(I_INT imax, T_P *pdf_old_gpu, I_INT pitch, int ntime, T_P *Boundary_xz0_gpu, T_P *Boundary_xz1_gpu);
extern __global__ void Out_ABC_xz(I_INT imax, T_P *pdf_old_gpu, I_INT pitch, int ntime, T_P *Boundary_xz0_gpu, T_P *Boundary_xz1_gpu);